//test_hip_kernel.cpp
#include "hip/hip_runtime.h"

#ifdef __cplusplus
extern "C" {
#endif

// 矩阵转置内核函数（添加了返回类型检查）
__global__ void matrixTranspose(float* out, float* in, const int width) {
    int x = blockDim.x * blockIdx.x + threadIdx.x;
    int y = blockDim.y * blockIdx.y + threadIdx.y;

    if (x < width && y < width) {
        out[y * width + x] = in[x * width + y];
    }
    // 虽然函数是void，但为了消除警告，这里可以添加空返回
    return;
}

// 启动内核的包装函数（修正了hipLaunchKernelGGL的使用）
void launch_matrix_transpose_kernel(
    void* d_output, void* d_input, int width,
    int grid_size_x, int grid_size_y,
    int threads_per_block_x, int threads_per_block_y,
    size_t shared_mem, size_t stream_ptr, int* hip_err
) {
    // 正确使用hipLaunchKernelGGL宏
    hipError_t err;
    hipStream_t stream = reinterpret_cast<hipStream_t>(stream_ptr); 
    //float* out = (float*)d_output;
    //float* in = (float*)d_input;
    

    //matrixTranspose <<<dim3(width / threads_per_block_x, width / threads_per_block_y),dim3(threads_per_block_x, threads_per_block_y)>>>(d_output, d_input, width);
    void* Args1[]={&d_output, &d_input, &width};
    err = hipLaunchKernel(
        (const void *)matrixTranspose,              // 内核函数名
        dim3(grid_size_x, grid_size_y), // 网格维度
        dim3(threads_per_block_x, threads_per_block_y), // 块维度
        Args1,                   // 共享内存大小
        0,                       // 流
        stream       // 内核参数
    );
    
    *hip_err = (int)hipGetLastError();
    return; // 显式返回以消除警告
}

#ifdef __cplusplus
}
#endif